#include "kuaicu.h"

#define THREAD_IN_BLOCK 96

__global__ void add0(float a[], const float v1[], const float v2[]) {
	int id = THREAD_ID;
	a[id] = v1[id]+v2[id];
}

__global__ void add1(float a[], int n, const float v1[], const float v2[]) {
	int id = THREAD_ID;
	if (id < n) {
		a[id] = v1[id]+v2[id];
	}
}

__global__ void add2(float a[], int n, int inca,
	const float v1[], int inc1, const float v2[], int inc2) {
	int id = THREAD_ID;
	if (id < n) {
		a[id*inca] = v1[id*inc1]+v2[id*inc2];
	}
}


void test1(kuai::cuda::FloatArray& a, const kuai::cuda::FloatArray& v1, const kuai::cuda::FloatArray& v2) {
	assert (a.size() == v1.size());
	assert (a.size() == v2.size());
	int n = a.size();
	int nblock = (n+THREAD_IN_BLOCK-1)/THREAD_IN_BLOCK;
	add1<<<nblock, THREAD_IN_BLOCK>>>(a.c_pointer(), n, v1.c_pointer(), v2.c_pointer());
}


void test2(kuai::cuda::FloatArray& a, const kuai::cuda::FloatArray& v1, const kuai::cuda::FloatArray& v2) {
	assert (a.size() == v1.size());
	assert (a.size() == v2.size());
	int n = a.size();
	int nblock = n/THREAD_IN_BLOCK;
	add0<<<nblock, THREAD_IN_BLOCK>>>(a.c_pointer(), v1.c_pointer(), v2.c_pointer());
	int offset=nblock*THREAD_IN_BLOCK;
	if (offset<n) {
		add1<<<nblock, THREAD_IN_BLOCK>>>(a.at(offset), n-offset, v1.at(offset), v2.at(offset));
	}
}

void test3(kuai::cuda::FloatArray& a, const kuai::cuda::FloatArray& v1, const kuai::cuda::FloatArray& v2) {
	assert (a.size() == v1.size());
	assert (a.size() == v2.size());
	int n = a.size();
	int nblock = (n+THREAD_IN_BLOCK-1)/THREAD_IN_BLOCK;
	add1<<<nblock, THREAD_IN_BLOCK>>>(a.c_pointer(), n, v1.c_pointer(), v2.c_pointer());
}

